#include "LinearInterpolation.h"




__global__ void interpolate_kernel(float* device_in, float* device_out, int width, int height){
	int x = blockIdx.x * blockDim.x + threadIdx.x;
	int y = blockIdx.y * blockDim.y + threadIdx.y;

	int out_x = x * 2;
	int out_y = y * 2;
	
	device_out[out_x + out_y * (width * 2)] = device_in[x + y * width];
	
	device_out[(out_x + 1) + out_y * (width * 2)] = 0.0;
	device_out[out_x + (out_y + 1) * (width * 2)] = 0.0;

	if(device_in[x + y * width] > 50){	
		if(device_in[(x + 1) + y * width] > 50)
			device_out[(out_x + 1) + out_y * (width * 2)] = 
				(device_in[x + y * width] + device_in[(x + 1) + y * width]) / 2.0f;

		if(device_in[x + (y + 1) * width])
			device_out[out_x + (out_y + 1) * (width * 2)] =
				(device_in[x + y * width] + device_in[x + (y + 1) * width]) / 2.0f;
	} 
}





float* LinearInterpolation::interpolate(xn::DepthMetaData* in){
	for(int y = 0; y < h; y++)
		for(int x = 0; x < w; x++)
			in_host[x + y * w] = (*in)(x, y);
	
	cudaMemcpy(in_device, in_host, sizeof(float) * w * h, cudaMemcpyHostToDevice);

	interpolate_kernel<<<dim3(20, 20), dim3(32, 24)>>>
		(device_in, device_out, w, h);

	cudaMemcpy(out_host, out_device, sizeof(float) * w * 2 * h * 2, cudaMemcpyDeviceToHost);
	return out_host;
}